#561 needs to be merged first, since this PR has branches from fix-opalx-release-assignment-segfault. After this PR I expect ALL weird GPU release seg-faults in OPALX to be fixed.

OPALX-project/OPALX#426 already documents a bug where OPALX seg-faults only on GPU (tested with CUDA) in Release mode during the accumulateHalo step after the scatter() kernel. In the following I'll try to outline what the symptoms are, then what the actual bug is (because they seem a bit disconnected), and my proposed fix.

What's happening and where?

When ParticleAttribute::scatter is is called, it calls accumulateHalo() after the scatter kernel which in turn calls the same function on the HaloCells instance, see here. Inside HaloCells::accumulateHalo we call the templated exchangeBoundaries<lhs_plus_assign>(view, layout, HALO_TO_INTERNAL);. This function then calls pack to load boundaries that need to be communicated into a buffer. The exact kernel that loads the buffer (see here) is where the seg-fault happens.

However, the problem is not the actual content. I've checked every object and every variable for consistency (ranges, memory spaces, values, values of member variables of objects and so on), but everything is initialized correctly. I also tried every compute sanitizer tool and debugger, but without any usable output, meaning the bug does not happen on GPU, but on host during the actual kernel launch.

Now, a "circumvention" of the bug was simply adding either an empty kernel launch or a fence before the templated function call, i.e.:

template <typename T, unsigned Dim, class... ViewArgs>
void HaloCells<T, Dim, ViewArgs...>::accumulateHalo(view_type& view, Layout_t* layout) {
	Kokkos::fence();
	exchangeBoundaries<lhs_plus_assign>(view, layout, HALO_TO_INTERNAL);
}

However, putting this fence before the BareField::accumulateHalo call inside ParticleAttribute::scatter, i.e. the following, does not work:

template <typename T, class... Properties>
template <typename Field, class PT, typename policy_type>
void ParticleAttrib<T, Properties...>::scatter(
	Field& f, const ParticleAttrib<Vector<PT, Field::dim>, Properties...>& pp,
	policy_type iteration_policy, hash_type hash_array) const {
	...
	Kokkos::parallel_for(
		"ParticleAttrib::scatter", iteration_policy,
		KOKKOS_LAMBDA(const size_t idx) {
			...
			detail::scatterToField(std::make_index_sequence<1 << Field::dim>{}, view, wlo, whi,
								   args, val);
		});

	...
	Kokkos::fence(); // This does NOT fix the seg-fault
	f.accumulateHalo();
	...
}

This indicates that it is not actually a (e.g.) device synchronization issue, but has something to do with how the compiler translates the call stack. Otherwise it would not depend on the placement of the fence in this simple "linear function call chain". Note the fence position changes generated code and linker selection and needs to be at these specific places before HaloCells.

Disassembling the core dump

The easiest (and turns out correct) interpretation is that the compiler optimizes something away that's then called leading to a host-side seg-fault. In that case, we need to look at what the compiler produced, so I let AI sift through the core dumps and the OPALX binary. It failed deterministically at program counter 0. The symbolized thread stack is (this we already know!):

0x0
Kokkos::parallel_for<MDRangePolicy<Cuda, Rank<3>>, HaloCells::pack functor>()
HaloCells<double, 3>::pack()
HaloCells<double, 3>::exchangeBoundaries<lhs_plus_assign>()
BareField<double, 3>::accumulateHalo()
ParticleAttrib<double>::scatter()

The failing instructions in the Kokkos::parallel_for instantiation are:

ldr x1, [x19, #0x2c8]
blr x1

x19 + 0x2c8 has address 0x183ee20. The symbol at that address is NVCC's translation-unit-local copy helper for the extended lambda in HaloCells<double, 3>::pack():

__nv_hdl_helper<...HaloCells<double, 3>::pack...>::fp_copier

Now the important part: The OPALX executable contains three copies of this same fp_copier helper:

0x183ee20 = 0x0
0x18427b0 = 0x0
0x18440a8 = 0x6138f0  // manager<pack lambda>::do_copy

As I understand it, NVCC identified three equivalent template specializations in three translation units for the same kernel launch and optimized them by merging all three ("weak, linker-coalesced Kokkos::parallel_for; something very similar as described here"). But since the NVCC stores the private fp_copier helper as a pointer to this function, the compiler didn't recognize it as something different. However, the failing pack now called the first helper 0x183ee20 that was part of a instantiation that "remained uninitialized". The host-side parallel_for launch mechanism therefore calls a null function pointer before the CUDA kernel launch.

This is why compute-sanitizer reported no device memory error and why source layout, empty kernels, and fence placement could make the crash appear or disappear between builds: sometimes the correct fp_copier is linked or sometimes the compiler recognized the different translation units.

During compilation, NVCC determines that concrete code is needed for HaloCells<double, 3>::pack(...). It therefore generates one copy of the template machinery and one private lambda-helper state. The same process happens independently for all three translation units (below I mentioned why there are at all three different translation units):

PartBunch.cpp
    -> HaloCells<double, 3>::pack
    -> private fp_copier for PartBunch.cpp

FieldSolver.cpp
    -> HaloCells<double, 3>::pack
    -> private fp_copier for FieldSolver.cpp

BinnedFieldSolver.cpp
    -> HaloCells<double, 3>::pack
    -> private fp_copier for BinnedFieldSolver.cpp

As mentioned before and as evident from the binary itself: the first two "unit-local fp_copier" are "optimized away", but the BinnedFieldSolver.cpp translation unit then links/runs against the very first fp_copier of PartBunch (which in turn is now NULL...).

Why does the bug only appear in OPALX and not in IPPL?

Looking at e.g. the LandauDamping binary, we see the same translation unit only once. The problem is that this specific part of IPPL is included three times in OPALX:

1. First we use forward instantiation in PartBunch, i.e. template class PartBunch<double, 3>;.
2. Then we instantiate another unit through FieldSolver<double, 3>::initOpenSolver().
3. Finally, we do the same in the other solver template class BinnedFieldSolver<double, 3>;.

These are three distinct translation units that contain (through ParticleAttribute and then BareField and HaloCells) this pack() call leading to the issue mentioned in the previous section.

Interesting side note: I am pretty sure (but speculation of course...) that the old FFTOpenPoissonSolver::pack seg-fault was very similar. Then, this was unintentionally "fixed" by moving the pack helper into their own public namespace, i.e. FieldBufferOps.hpp. Why this would fix it is explained in the next section.

The fix

The fix replaces NVCC extended lambdas with named functor types. These store the Kokkos views directly and use default C++ copy construction. Therefore, the duplicate template instantiation can be safely merged by the linker without depending on the private (private meaning "per translation unit") fp_copier or fp_deleter helpers.

The original code needed these helpers, because the following pattern constructs the already mentioned unnamed extended lambda:

ippl::parallel_for(
    "HaloCells::pack()", policy,
    KOKKOS_LAMBDA(const index_array_type& args) {
        buffer(l) = apply(subview, args);
    });

...which needs these fp_copier and fp_deleter helpers.

unpack() and applyPeriodicSerialDim() use the same extended-lambda pattern and also need to be changed with pack() or they can seg-fault in a similar way.

Testing

I ran SwissFEL-booster-sc on 16 ranks as well as all regression tests in release on 2 GH200 ranks without any of the previous OPALX "bug circumventions". Everything works fine as expected.

Additional comments

I also included a small potentially uninitialized variable fix: BufferHandler.h has two members usedSize_m and freeSize_m which are in some instances not instantiated, but first accessed via usedSize_m += value (without any assignment first). This did not lead to bug, but could be dangerous.

Finally, I am pretty sure that what we do in OPALX (extended lambdas behind multiple translation units on heavily templated code) is completely legal, but exposed some kind of Compiler bug. Note that I don't know what the actual NVCC bug could be, but it looks like perhaps the compiler was simply overwhelmed.